\hypertarget{structnih_1_1cuda_1_1_kd__builder}{
\section{nih\-:\-:cuda\-:\-:\-Kd\-\_\-builder$<$ \-Integer $>$ \-Struct \-Template \-Reference}
\label{structnih_1_1cuda_1_1_kd__builder}\index{nih\-::cuda\-::\-Kd\-\_\-builder$<$ Integer $>$@{nih\-::cuda\-::\-Kd\-\_\-builder$<$ Integer $>$}}
}


{\ttfamily \#include $<$kd\-\_\-builder.\-h$>$}

\subsection*{\-Public \-Member \-Functions}
\begin{DoxyCompactItemize}
\item 
{\footnotesize template$<$typename Output\-Tree , typename Iterator $>$ }\\void \hyperlink{structnih_1_1cuda_1_1_kd__builder_a227e02ff324dfb9c251fee6954e942c8}{build} (\-Output\-Tree \&out\-\_\-tree, thrust\-::device\-\_\-vector$<$ uint32 $>$ \&out\-\_\-index, const \hyperlink{structnih_1_1_bbox}{\-Bbox3f} bbox, const \-Iterator points\-\_\-begin, const \-Iterator points\-\_\-end, const uint32 max\-\_\-leaf\-\_\-size)
\end{DoxyCompactItemize}
\subsection*{\-Public \-Attributes}
\begin{DoxyCompactItemize}
\item 
\hypertarget{structnih_1_1cuda_1_1_kd__builder_aa620524c7d26050cd6713b2fd619ed45}{
thrust\-::device\-\_\-vector$<$ \-Integer $>$ {\bfseries m\-\_\-codes}}
\label{structnih_1_1cuda_1_1_kd__builder_aa620524c7d26050cd6713b2fd619ed45}

\item 
\hypertarget{structnih_1_1cuda_1_1_kd__builder_aadd42d7e987727d8d480debfecc8e2fe}{
uint32 {\bfseries m\-\_\-levels} \mbox{[}64\mbox{]}}
\label{structnih_1_1cuda_1_1_kd__builder_aadd42d7e987727d8d480debfecc8e2fe}

\item 
\hypertarget{structnih_1_1cuda_1_1_kd__builder_a96b9aebcce9a59a0c9bfe82d973ce547}{
\hyperlink{structnih_1_1_bbox}{\-Bbox3f} {\bfseries m\-\_\-bbox}}
\label{structnih_1_1cuda_1_1_kd__builder_a96b9aebcce9a59a0c9bfe82d973ce547}

\item 
\hypertarget{structnih_1_1cuda_1_1_kd__builder_adfbb54bfad7dc7176f6d87585a1ab29c}{
uint32 {\bfseries m\-\_\-node\-\_\-count}}
\label{structnih_1_1cuda_1_1_kd__builder_adfbb54bfad7dc7176f6d87585a1ab29c}

\item 
\hypertarget{structnih_1_1cuda_1_1_kd__builder_a43a932652753fed10014c39973e5ac2c}{
uint32 {\bfseries m\-\_\-leaf\-\_\-count}}
\label{structnih_1_1cuda_1_1_kd__builder_a43a932652753fed10014c39973e5ac2c}

\item 
\hypertarget{structnih_1_1cuda_1_1_kd__builder_a158d2483ee18b23b53c9b1b5e9e55106}{
\hyperlink{structnih_1_1cuda_1_1_bintree__gen__context}{cuda\-::\-Bintree\-\_\-gen\-\_\-context} {\bfseries m\-\_\-kd\-\_\-context}}
\label{structnih_1_1cuda_1_1_kd__builder_a158d2483ee18b23b53c9b1b5e9e55106}

\end{DoxyCompactItemize}


\subsection{\-Detailed \-Description}
\subsubsection*{template$<$typename Integer$>$struct nih\-::cuda\-::\-Kd\-\_\-builder$<$ Integer $>$}

\-G\-P\-U-\/based middle-\/split k-\/d tree builder

\-This class provides the context to generate k-\/d trees on the \-G\-P\-U starting from a set of unordered points. \-The output is a set of nodes with the corresponding leaves and a set of primitive indices into the input set of points. \-The output leaves will specify contiguous ranges into this index.


\begin{DoxyTemplParams}{\-Template Parameters}
{\em \-Integer} & an integer type that determines the number of bits used to compute the points' \-Morton codes. \-Accepted values are uint32 and uint64.\\
\hline
{\em \-Output\-Tree} & a template class used to write the output tree, with the following interface\-:\\
\hline
\end{DoxyTemplParams}

\begin{DoxyCode}
 struct OutputTree
 {
    void reserve_nodes(const uint32 n);  // reserve space for n nodes
    void reserve_leaves(const uint32 n); // reserve space for n leaves

    Context get_context();             // get a context to write nodes/leaves

    struct Context
    {
        void write_node(
           const uint32 node,          // node to write
           const uint32 offset,        // child offset
           const uint32 skip_node,     // skip node
           const uint32 begin,         // node range begin
           const uint32 end,           // node range end
           const uint32 split_index,   // split index
           const uint32 split_dim,     // splitting dimension
           const uint32 split_plane);  // splitting plane

        void write_node(
           const uint32 node,          // node to write
           const uint32 offset,        // child offset
           const uint32 skip_node,     // skip node
           const uint32 begin,         // node range begin
           const uint32 end);          // node range end

        void write_leaf(
           const uint32 index,         // leaf to write
           const uint32 begin,         // leaf range begin
           const uint32 end);          // leaf range end
    };
 };
\end{DoxyCode}


\-The following code snippet shows how to use this builder\-:


\begin{DoxyCode}
 #include <nih/kd/cuda/kd_builder.h>
 #include <nih/kd/cuda/kd_context.h>

 thrust::device_vector<Vector3f> points;
 ... // code to fill the input vector of points

 thrust::device_vector<Kd_node>  kd_nodes;
 thrust::device_vector<uint2>    kd_leaves;
 thrust::device_vector<uint32>   kd_index;

 nih::cuda::Kd_builder<uint64> builder( kd_index );
 nih::cuda::Kd_context kd_tree( &kd_nodes, &kd_leaves, NULL );
 builder.build(
     kd_tree,                                    // output tree
     kd_index,                                   // output index
     Bbox3f( Vector3f(0.0f), Vector3f(1.0f) ),   // suppose all bboxes are in
       [0,1]^3
     points.begin(),                             // begin iterator
     points.end(),                               // end iterator
     4 );                                        // target 4 objects per leaf
\end{DoxyCode}
 

\subsection{\-Member \-Function \-Documentation}
\hypertarget{structnih_1_1cuda_1_1_kd__builder_a227e02ff324dfb9c251fee6954e942c8}{
\index{nih\-::cuda\-::\-Kd\-\_\-builder@{nih\-::cuda\-::\-Kd\-\_\-builder}!build@{build}}
\index{build@{build}!nih::cuda::Kd_builder@{nih\-::cuda\-::\-Kd\-\_\-builder}}
\subsubsection[{build}]{\setlength{\rightskip}{0pt plus 5cm}template$<$typename Integer $>$ template$<$typename Output\-Tree , typename Iterator $>$ void {\bf nih\-::cuda\-::\-Kd\-\_\-builder}$<$ \-Integer $>$\-::build (
\begin{DoxyParamCaption}
\item[{\-Output\-Tree \&}]{out\-\_\-tree, }
\item[{thrust\-::device\-\_\-vector$<$ uint32 $>$ \&}]{out\-\_\-index, }
\item[{const {\bf \-Bbox3f}}]{bbox, }
\item[{const \-Iterator}]{points\-\_\-begin, }
\item[{const \-Iterator}]{points\-\_\-end, }
\item[{const uint32}]{max\-\_\-leaf\-\_\-size}
\end{DoxyParamCaption}
)}}
\label{structnih_1_1cuda_1_1_kd__builder_a227e02ff324dfb9c251fee6954e942c8}
build a bvh given a set of points that will be reordered in-\/place


\begin{DoxyParams}{\-Parameters}
{\em out\-\_\-tree} & output tree \\
\hline
{\em out\-\_\-index} & output index \\
\hline
{\em bbox} & global bbox \\
\hline
{\em points\-\_\-begin} & beginning of the point sequence to sort \\
\hline
{\em points\-\_\-end} & end of the point sequence to sort \\
\hline
{\em max\-\_\-leaf\-\_\-size} & maximum leaf size \\
\hline
\end{DoxyParams}


\-The documentation for this struct was generated from the following files\-:\begin{DoxyCompactItemize}
\item 
kd/cuda/\hyperlink{kd__builder_8h}{kd\-\_\-builder.\-h}\item 
kd/cuda/kd\-\_\-builder\-\_\-inline.\-h\end{DoxyCompactItemize}
